[Tools] Add architecture-general ISA resource diff tool - #1015

gfx1250 is based on cdna5 arch.

Done
You're right, and thanks — the label was wrong. Fixed in 8ef12c8.
Rather than relabel it cdna5, I dropped the microarchitecture names from this code entirely, because the parser never branches on the product family. It branches on two ISA properties.

[blocking] This split only recognizes the abbreviated amdgcn-amd-amdhsa--gfx950 form used by the synthetic test. LLVM also emits target IDs such as amdgcn-amd-amdhsa-unknown-gfx950; for those, parse_target_id() returns processor=None. Because compare_snapshots() rejects an architecture mismatch only when both processors are known, an unknown-gfx942 snapshot compared with unknown-gfx950 currently returns exit 0 / RESULT: OK. Please extract the trailing concrete gfx... processor from both LLVM spellings (and fail closed when it is genuinely unavailable), with a regression test proving a 942/950 mismatch exits 2.

Done in b491815e. Processor now comes from the last - separated field, so both spellings parse.

The guard needed fixing too, otherwise the parse fix alone would not have closed it: with both processors None, processor != processor is false, so it never reached the known check. It now blocks whenever a processor is unidentified and the target IDs differ.

Tests: both spellings, 942 vs 950 → exit 2, gfx11-generic vs gfx12-generic → exit 2.

[blocking] A discovered *final_isa.s that yields no records is omitted as a warning, so the comparison can become silently partial. I reproduced this with one valid dump plus one empty/truncated final_isa.s on each side: both bad files are skipped and diff returns exit 0 / RESULT: OK for the remaining kernel. That contradicts the stated fail-closed contract. Please record this as a problem (or retain a blocked sentinel entry) so any final-ISA file the tool cannot parse forces exit 2, and cover the mixed valid/invalid tree case.

Done in b491815e. Recorded as a problem instead of a warning, so it reaches exit 2 — and summarize/snapshot fail closed on it too, via Snapshot.trustworthy.

Took the problem over a sentinel entry: a sentinel on both sides yields one blocked record per metric, and on one side it degrades into ONLY IN AFTER.

Test covers the mixed valid/invalid tree, with the healthy half asserted to exit 0 on its own first.

This points readers to a Compare per-kernel ISA resources section that is not present in docs/testing_benchmarking_guide.md; this PR only adds the script to that guide's source-file table. Please either add the promised column reference or link to the documentation that actually defines these columns.

Done in b491815e. Dropped the dangling reference and put the column table in the skill itself.

docs/kernel_tuning_guide.md:462 already pointed here for the full column reference, so the two files pointed at each other while it existed in neither. Defining it here makes that pointer true, rather than adding an ISA column reference to a testing guide.

[blocking] Processor equality does not prove the target modes are comparable. I changed only the suffixes to gfx942:xnack+ and gfx942:xnack-; this branch returned exit 0 / RESULT: OK. Removing both target declarations also returns exit 0. These modes can change code generation and resource use, so please normalize the empty-vs-unknown environment spelling but require a nonempty processor and identical feature sets, otherwise exit 2. Add regressions for differing features and two missing targets.

Done in aa9b7307. Comparability is now decided on the parsed processor and feature set instead of the raw target ID string, so the environment field normalizes away and xnack/sramecc no longer slip through. Features are sorted on parse, since LLVM's ordering is not a contract.

Both halves confirmed: gfx942:xnack+ vs gfx942:xnack- → exit 2, xnack+ vs sramecc+:xnack+ → exit 2, no target directive on either side → exit 2, and amdgcn-amd-amdhsa--gfx942 vs amdgcn-amd-amdhsa-unknown-gfx942 stays exit 0.

One consequence worth naming: requiring a nonempty processor also refuses gfx11-generic against itself, since RE_PROC does not recognize it. That is the right default — without a processor the scratch/spill applicability is undecidable — and if generic targets ever matter here the fix is to teach RE_PROC about them, not to loosen the verdict. Pinned in the tests either way.

[blocking] This silently drops a metadata entry with neither .name nor .symbol; the resulting snapshot still has trustworthy=True. records[name] = ... also lets a later duplicate name overwrite an earlier record, so a stale duplicate can hide a changed counter. Treat missing and duplicate kernel identities as snapshot problems that force exit 2, and cover both mutations.

Done in aa9b7307. Both are now snapshot problems.

The identity-less case was a plain bug: _kernel_name() already returned "metadata entry has neither .name nor .symbol" and the caller discarded it before continue. parse_isa() now returns (records, file problems) so there is a place for a fault that belongs to the file rather than to a record; collect() folds them into Snapshot.problems.

Duplicates keep the first entry rather than the last and record the collision, so a stale duplicate can no longer overwrite a changed counter silently.

Both regressions splice a second metadata entry into an otherwise healthy file — a single-entry file already raises on losing its only kernel, which hides the case you reported.

[blocking] Resource counters are non-negative, but int(raw) accepts impossible values. Changing .vgpr_count from 285 to -1 was classified as an improvement and returned exit 0. Validate the metric domain and make any negative resource count untrustworthy instead of comparable.

Done in aa9b7307. Negative values are rejected in both _metadata_cell() and _symbol_cell(), which makes the metric unparsed and forces exit 2.

Agreed this was the worst of the five: every metric here is a count or a byte size, so a negative one means the dump is malformed, and rendering the drop toward it as an improvement is the one verdict the tool must never invent. Regression covers .vgpr_count: -1 and a negative .set symbol.

[blocking] Replacement decoding can silently erase an instruction mnemonic. Corrupting one byte in one ds_read changed the reported count from 3 to 2 while the snapshot remained trustworthy with no problems. Decode strictly, or detect replacement characters and force exit 2, rather than publishing partial instruction counts.

Done in aa9b7307. Kept the lenient decode — a traceback reports worse than exit 2 does — but the file is now decoded strictly first, and a UnicodeDecodeError is recorded as a file problem before falling back to errors="replace". So the replacement is admitted rather than counted.

Went with strict-then-fall-back over scanning for U+FFFD, which would also fire on a file that legitimately contains that character. Regression corrupts one byte inside a ds_read mnemonic and asserts both halves: the count really does drop to 2, and the snapshot is no longer trustworthy.

[blocking] Schema v2 accepts n/a for every metric. Setting mandatory vgpr to n/a on both JSON inputs bypasses comparison and returns exit 0. Validate state applicability by metric/architecture: mandatory register, spill, scratch-byte, and static-LDS counters cannot be n/a; reserve it only for quantities that genuinely do not exist on that target.

Declining this one, unlike the other four.

n/a is only ever produced by _instruction_cells(), for scratch_store/scratch_load on a target that spills through buffer_*. No path in this tool — or in LLVM — emits n/a for vgpr, and Cell.from_json() already validates the state enum, the value type and the reason type. Reaching the case you describe means hand-editing the tool's own serialized output, which is not a trustworthiness gap in the data path the way the other four are: those all start from a file LLVM or the filesystem produced.

The prescribed fix also has a cost I would rather not pay. "Validate state applicability by metric/architecture" puts a second copy of the applicability rules in the loader, so spills_via_scratch() and the JSON validator would have to be kept in agreement forever, and the first time they drift the tool refuses a snapshot it wrote itself.

Happy to reconsider if you can show a snapshot the tool actually produces that carries n/a on a mandatory metric, or a second producer of this schema — either would move it into the same class as the rest.

[blocking] The assembly parser now rejects negative counters, but the reusable JSON decoder still accepts them. I reproduced a schema-v2 snapshot with vgpr=-1: load_snapshot() marked it trustworthy and diffing 10 -> -1 returned exit 0 / RESULT: OK. Please reject negative VALUE cells during snapshot loading (and add a CLI-level JSON regression). This is the same fail-closed boundary as the still-unresolved mandatory-n/a case, where two snapshots both declaring vgpr as n/a also return OK.

Split verdict on the two halves.

Negative VALUE cells on load — accepted, fixed in 7eff378a. You are right that this was an asymmetry I introduced: I put the domain check in _metadata_cell()/_symbol_cell() and left Cell.from_json() without it, for the same Cell type and the same invariant. Non-negativity belongs to the metric, not to the parser that happened to read it. Cell.from_json() now raises SnapshotError on a negative value, consistent with how it already rejects a wrong type there.

CLI-level regression added as asked: diff ten.json negative.json → exit 2, with a healthy JSON pair asserted at exit 0 first so the 2 is attributable. Mutation-checked — disabling the new branch fails exactly that test.

Mandatory n/a — still declining, for the reason given last round, which this review does not address. The distinction is domain versus applicability. Non-negativity is intrinsic to the value and is one comparison in the place the value is constructed, which is why I took it. Which metrics may be n/a is a property of the target, derived by _instruction_cells() from spills_via_scratch(); asserting it in the loader means a second copy of that rule, and the first time the two drift the tool refuses a snapshot it wrote itself.

I asked for one of two things to move it into the same class as the rest: a snapshot the tool actually produces that carries n/a on a mandatory metric, or a second producer of this schema. Re-running the same hand-edited input is not either of those — Cell.na() is called in exactly one place, for scratch_store/scratch_load, and no path in this tool or in LLVM writes n/a for vgpr. If you have a concrete producer in mind I will take it.

Unrelated but worth flagging: this review was against a3e74b4, which was a force-push of the same commit onto a newer main from another machine. I rebased onto it rather than over it, so nothing from that push was dropped — the only delta from a3e74b4 is the fix above.

Closing my half of this thread: the mandatory-n/a finding is withdrawn.

I checked it the way I should have in round 3 instead of restating it. Cell.na has exactly two construction sites in the file, lines 540-541, both scratch_store/scratch_load under spills_via_scratch() is False. There is no third, so no dump LLVM emits and no path through parse_isa can put n/a on a mandatory metric -- reaching it means editing the tool's own serialized output by hand. That is the distinction you drew between this and the other four, and it holds.

The negative-VALUE half is verified fixed on 7eff378a, independently of your test: healthy JSON pair -> exit 0, 10 -> -1 -> exit 2, with the error naming the cell (k::gemm_0.vgpr.value is negative: -1).

[blocking] vgpr is the only VGPR-family trigger, and that is sound only where .vgpr_count is the arch+accumulator sum. LLVM makes it the sum only on gfx90a+:

// AMDGPUMCExpr::evaluateTotalNumVGPR
bool Has90AInsts = AMDGPU::isGFX90A(STI);
uint64_t TotalNum = Has90AInsts && NumAGPR ? alignTo(NumVGPR, 4) + NumAGPR
                                           : std::max(NumVGPR, NumAGPR);

isGFX90A tests FeatureGFX90AInsts. gfx908 (FeatureISAVersion9_0_8) is MFMA-capable -- it emits .agpr_count, which is gated on STM.hasMAIInsts() -- but does not carry that feature. So on gfx908 .vgpr_count is max(arch, agpr), not the sum.

Reproduced on head 7eff378a with a gfx908 dump whose .vgpr_count is exactly what LLVM would emit: arch_vgpr 64 on both sides, agpr 10 -> 60, so max(64,10) == max(64,60) == 64.

kernel    *vgpr arch_vgpr        agpr *sgpr ... *vgpr_spill ...
k::gemm_0    64        64 10->60(+50)    59 ...           0 ...

compared 1 of 1 kernels; 0 unchanged; 1 changed; worsened: 0; improved: 0
RESULT: OK          # exit 0

Fifty accumulator registers on a part whose 256-entry AGPR file is a separate occupancy limiter, reported as no regression. The tool has already parsed the number and prints it in the row -- it simply is not a trigger, and agpr is INFO.

docs/kernel_tuning_guide.md in this same PR states the precondition correctly ("targets with a unified register file (gfx90a and later MFMA-capable parts)"). The code applies the conclusion unconditionally, so the one target where the premise fails is the one target where the premise was worth checking.

Please make it decidable from the target rather than assumed -- either promote agpr to a trigger when the target is MFMA-capable but not unified-RF, or make vgpr unparsed (exit 2) when vgpr < arch_vgpr + agpr on a target this tool cannot prove unified. A cross-check is the cheaper of the two and matches what _cross_check_scratch already does for the other quantity that has two sources.

Worth a regression shape either way: TARGETS covers unified-RF-with-AGPRs (gfx942) and no-AGPR (gfx1250), but not has-AGPRs-without-unified-RF, which is exactly where this lives.

[non-blocking] parse_target_id sorts the feature set "because the feature set is what matters and LLVM's order is not a contract", and Arch's own field comment states the invariant ("sorted, so two spellings of one set compare equal"). Arch.from_json does not apply it -- it takes tuple(features) verbatim, so the invariant holds on the assembly path and not on the snapshot path:

parse_target_id("...gfx942:xnack+:sramecc+").features  -> ('sramecc+', 'xnack+')
Arch.from_json({... "features": ["sramecc+","xnack+"]}) -> ('sramecc+', 'xnack+')
Arch.from_json({... "features": ["xnack+","sramecc+"]}) -> ('xnack+', 'sramecc+')

_target_blockers(first, second)
-> ('target features differ (sramecc+:xnack+ vs xnack+:sramecc+); they change code
    generation, so this diff would measure the build flags rather than the change
    under test',)

Fail-closed, so no wrong verdict -- but it is a false exit 2 between two snapshots the tool would call identical had it re-parsed the target ID, and it is the same class of "one normalization, applied on one of the two paths that build the object" as the negative-value asymmetry you fixed in 7eff378a. tuple(sorted(features)) here is the call parse_target_id already makes.

[non-blocking] Last-wins here, unlike the metadata duplicate path you changed to first-wins-and-record in aa9b7307. Records are keyed by base after the .L prefix is stripped, so the two spellings of one name collide and the later line silently replaces the earlier one with nothing recorded:

.set gemm_0.num_vgpr, 256
.set .Lgemm_0.num_vgpr, 1

-> arch_vgpr = Cell(state='value', value=1)
   file problems = ()

Reachability is low -- it needs one file emitting both spellings for one base name -- but the comment two lines up is what makes it worth closing: both spellings are in circulation (90 of 233 real dumps use the .L form), and this normalization is what puts them in the same key space. The resolution you already picked for duplicate metadata entries applies unchanged: keep the first, record the collision.

private_seg_size is the one with teeth. It is not a displayed metric, it feeds _cross_check_scratch, so a collided value produces a scratch-bytes disagreement that does not exist and takes the run to exit 2 with a wrong reason.

[non-blocking] load_snapshot restores record.problems faithfully at line 873 and then nothing ever reads them. Snapshot.problems gets one generic line instead, so a snapshot carrying a real per-kernel fault round-trips into a comparison strictly less diagnosable than the directory it came from:

record.problems restored:  ("metadata .name 'gemm_0' disagrees with .symbol 'other_name.kd'",)
snapshot.problems surfaced: ('bad.json was written from data the tool could not fully parse',)

The verdict is right (exit 2) and the detail is sitting in the file, which is what makes losing it avoidable: diff dirA dirB names the fault, diff a.json b.json does not, for the same data. Folding the restored per-record problems into Snapshot.problems -- what collect() already does at line 745 -- makes the two input paths report the same thing.

[non-blocking] This is compatibility code for a format the repository has never shipped. scripts/isa_resource_table.py is new in this PR, SCHEMA_VERSION has been 2 since 3b3e9e8b, and aa6a2d40 is the only commit that ever wrote a v1 snapshot -- it never left this branch, so no v1 file exists outside a working tree that also contains the tool that reads v2.

It also cannot produce a comparable row even if one did. _load_v1 builds KernelRecord(..., arch=Arch()), which _target_blockers refuses as an unidentified target, and the five metrics with no v1 alias come back Cell.unparsed. Diffing a v1 snapshot against itself:

gemm_0: unparsed metric numbered_sgpr: absent from this v1 snapshot
gemm_0: unparsed metric lds_write: absent from this v1 snapshot
gemm_0: unparsed metric matrix_ops: absent from this v1 snapshot
5 kernel(s)/metric(s) not trustworthy (see stderr)
RESULT: NOT TRUSTWORTHY      # exit 2

So the branch, the alias table and the warning are ~35 lines that can only ever reach the verdict the schema check three lines below would reach on its own. Suggest deleting it. If you would rather keep it, the warning should say the snapshot cannot be compared at all rather than "regenerate it with this version of the tool", which reads as though the current run produced a usable answer.

[non-blocking] Both stdout-shape claims are false exactly when there are problems -- the case an agent scripting against this contract most needs to get right. _render_problems prints its count of untrustworthy items to stdout, between the count line and the verdict:

$ isa_resource_table.py diff <two real gfx950 dump trees> 2>/dev/null | tail -3
compared 2 of 6 kernels; 2 unchanged; 0 changed; worsened: 0; improved: 0
4 kernel(s)/metric(s) not trustworthy (see stderr)
RESULT: NOT TRUSTWORTHY

The last two stdout lines are that count and the verdict, not "a count line and a RESULT: verdict". Line 147's "-q/--quiet prints only the verdict line" fails the same way, on the same run:

$ isa_resource_table.py diff <same two trees> -q 2>/dev/null
4 kernel(s)/metric(s) not trustworthy (see stderr)
RESULT: NOT TRUSTWORTHY

Two ways to close it. Narrow both sentences to "the last stdout line is always the RESULT: verdict" -- true unconditionally, and all a script needs. Or route that one line to stderr with the rest of _render_problems, which would make both sentences true as written and match "Warnings on stderr never change the exit code" in Pitfalls.

[non-blocking] Same boundary as the blocking vgpr finding in round 4. "The accumulator half of that total" holds only where .vgpr_count is the arch+accumulator sum, which LLVM makes it only on gfx90a+; on gfx908 it writes max(arch, agpr), so agpr is not a component of vgpr at all there.

That makes the "Do not add arch_vgpr and agpr to vgpr" note below read as a stronger guarantee than the metadata format gives, and it is the sentence a reader will use to decide that a moved AGPR count needs no attention.

docs/kernel_tuning_guide.md in this same PR already scopes it correctly ("targets with a unified register file (gfx90a and later MFMA-capable parts)"). Carrying that same qualifier into this table -- and into the vgpr row above it -- would make the two documents agree. Whatever the code ends up doing about the trigger, this table should describe the format, not the gfx90a+ special case.

[non-blocking] matrix_ops misses v_smfmac_*, CDNA3/4's sparse MFMA, so a kernel built on the sparse matrix cores reports matrix_ops 0. On head:

_categorize("v_smfmac_f32_16x16x32_f16")  -> None
_categorize("v_mfma_f32_16x16x16_f16")    -> matrix_ops

Informational column, so no verdict is wrong; the fix is ^v_(?:mfma|wmma|smfmac). Flagging it from the ISA rather than from a failure I hit -- none of the 61 real dumps on this machine use sparse MFMA -- so treat it as a completeness note on the same footing as the ds_read/ds_load spelling pair this tuple already handles. The SKILL.md row for this column says "MFMA / WMMA sites" and would want the same edit.

[non-blocking] TARGETS covers both sides of three axes -- .agpr_count present or absent, gfx9 or gfx11+ LDS spelling, empty or unknown environment -- and the file's docstring says so. The axis it does not cover is the one that carries the round-4 blocking finding: whether the target's register file is unified.

The two shapes are gfx942 (unified RF, AGPRs in use, vgpr_count = 256 + 29) and gfx1250 (no AGPRs at all). Both satisfy ".vgpr_count is the arch+accumulator total". gfx908 is the third case -- MFMA-capable, .agpr_count emitted, but no FeatureGFX90AInsts, so LLVM writes max(arch, agpr) -- and on that shape the only VGPR trigger cannot see an AGPR-only change.

A gfx908 entry would be a two-line addition here (arch="gfx908", mfma=True, vgpr_count=max(num_vgpr, num_agpr)) and would fail today against any fix that makes the AGPR change visible, which is what makes it worth pinning rather than describing.